#include <cuda_runtime.h>
#include <iostream>
/*
对输入的每个像素点进行插值计算。它通过将输入的 tenIn 和光流 tenFlow 转换为目标输出的像素值分布。以下是插值计算的核心逻辑：

    根据光流值，计算像素的目标位置 (fltX, fltY)。
    根据目标位置 (fltX, fltY)，找到其周围的四个像素点（北西、北东、南西、南东）的坐标。
    计算四个像素点对目标位置的贡献权重（双线性插值）。
    使用 atomicAdd 将插值后的值累加到对应位置。
*/
extern "C" __global__ void __launch_bounds__(512) softsplat_out(
    const int n,
    const float* __restrict__ tenIn,
    const float* __restrict__ tenFlow,
    float* __restrict__ tenOut
) { for (int intIndex = (blockIdx.x * blockDim.x) + threadIdx.x; intIndex < n; intIndex += blockDim.x * gridDim.x) {
    const int intN = ( intIndex / SIZE_3(tenOut) / SIZE_2(tenOut) / SIZE_1(tenOut) ) % SIZE_0(tenOut);
    const int intC = ( intIndex / SIZE_3(tenOut) / SIZE_2(tenOut)                  ) % SIZE_1(tenOut);
    const int intY = ( intIndex / SIZE_3(tenOut)                                   ) % SIZE_2(tenOut);
    const int intX = ( intIndex                                                    ) % SIZE_3(tenOut);

    assert(SIZE_1(tenFlow) == 2);

    float fltX = (float) (intX) + VALUE_4(tenFlow, intN, 0, intY, intX);
    float fltY = (float) (intY) + VALUE_4(tenFlow, intN, 1, intY, intX);

    if (isfinite(fltX) == false) { return; }
    if (isfinite(fltY) == false) { return; }

    float fltIn = VALUE_4(tenIn, intN, intC, intY, intX);

    int intNorthwestX = (int) (floor(fltX));
    int intNorthwestY = (int) (floor(fltY));
    int intNortheastX = intNorthwestX + 1;
    int intNortheastY = intNorthwestY;
    int intSouthwestX = intNorthwestX;
    int intSouthwestY = intNorthwestY + 1;
    int intSoutheastX = intNorthwestX + 1;
    int intSoutheastY = intNorthwestY + 1;

    float fltNorthwest = ((float) (intSoutheastX) - fltX) * ((float) (intSoutheastY) - fltY);
    float fltNortheast = (fltX - (float) (intSouthwestX)) * ((float) (intSouthwestY) - fltY);
    float fltSouthwest = ((float) (intNortheastX) - fltX) * (fltY - (float) (intNortheastY));
    float fltSoutheast = (fltX - (float) (intNorthwestX)) * (fltY - (float) (intNorthwestY));

    if ((intNorthwestX >= 0) && (intNorthwestX < SIZE_3(tenOut)) && (intNorthwestY >= 0) && (intNorthwestY < SIZE_2(tenOut))) {
        atomicAdd(&tenOut[OFFSET_4(tenOut, intN, intC, intNorthwestY, intNorthwestX)], fltIn * fltNorthwest);
    }

    if ((intNortheastX >= 0) && (intNortheastX < SIZE_3(tenOut)) && (intNortheastY >= 0) && (intNortheastY < SIZE_2(tenOut))) {
        atomicAdd(&tenOut[OFFSET_4(tenOut, intN, intC, intNortheastY, intNortheastX)], fltIn * fltNortheast);
    }

    if ((intSouthwestX >= 0) && (intSouthwestX < SIZE_3(tenOut)) && (intSouthwestY >= 0) && (intSouthwestY < SIZE_2(tenOut))) {
        atomicAdd(&tenOut[OFFSET_4(tenOut, intN, intC, intSouthwestY, intSouthwestX)], fltIn * fltSouthwest);
    }

    if ((intSoutheastX >= 0) && (intSoutheastX < SIZE_3(tenOut)) && (intSoutheastY >= 0) && (intSoutheastY < SIZE_2(tenOut))) {
        atomicAdd(&tenOut[OFFSET_4(tenOut, intN, intC, intSoutheastY, intSoutheastX)], fltIn * fltSoutheast);
    }
} }